#ifndef NV_KERNEL
#include <hip/hip_runtime.h>
#endif

extern "C" __global__ void Hello(float * c, const float * a, const unsigned char * mask, 
	const double bias,const unsigned int length)
{
#ifdef NV_KERNEL
	unsigned int loc_id_x = threadIdx.x;
	unsigned int glb_id_x = blockDim.x * blockIdx.x + threadIdx.x;
	unsigned int glb_id_y = blockDim.y * blockIdx.y + threadIdx.y;
#else
	unsigned int loc_id_x = hipThreadIdx_x;
	unsigned int glb_id_x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
	unsigned int glb_id_y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
#endif

	if (glb_id_x >= length)
		return;

	if (glb_id_x < length / 2)
	{
		if (mask[glb_id_x] == 0)
			c[glb_id_x] = a[glb_id_x] - (float)bias;
		else
			c[glb_id_x] = a[glb_id_x] + (float)bias;
	}
	else
	{
		c[glb_id_x] = 3.1415926f;
	}
}
